#ifndef GPU_DART_H
#define GPU_DART_H

#include <cuda.h>
#include <curand_kernel.h>
#include "vectorGpu.h"

class GPUDart
{
public:
	struct GPUMDVector
	{
		void init(unsigned int size, unsigned int dimensions) { size_ = size; GPUMALLOC(&p_, size * dimensions * sizeof(DATA_TYPE)); }
		void destroy() { GPUFREE(p_); }

		__device__ DATA_TYPE& operator()(unsigned int index, unsigned int d) { return p_[d * size_ + index]; }
		__device__ DATA_TYPE operator()(unsigned int index, unsigned int d) const { return p_[d * size_ + index]; }

		DATA_TYPE* p_;
		unsigned int size_;
	};

	struct GPUInterval
	{
		DATA_TYPE intervalBegins[MAX_INTERVAL];
		DATA_TYPE intervalEnds[MAX_INTERVAL];
		short next[MAX_INTERVAL];
		short prev[MAX_INTERVAL];

		short head;
		short tail;
		short numIntervals;
		short numUsed;

		__device__ void insert(short pos, DATA_TYPE begin, DATA_TYPE end);
		__device__ void erase(short from, short to);

		static __device__ short advance(short* indexArray, short it, short step);
		static __device__ short lower_bound(DATA_TYPE* dataArray, short* indexArray, short first, short last, DATA_TYPE val, short* distance);
		static __device__ short upper_bound(DATA_TYPE* dataArray, short* indexArray, short first, short last, DATA_TYPE val, short* distance);
	};

	GPUDart() : useMultipleBlocks_(false), validateResults_(true), h_numSamples_(0), h_failureCount_(0), h_m_(0) {}
	virtual ~GPUDart() {}

	virtual bool init(unsigned int m, DATA_TYPE radius, unsigned int d) = 0;
	virtual void setupDarts() = 0;
	virtual void updateIntervals() = 0;
	virtual void updateIntervals2() = 0;
	virtual void pickPoint() = 0;
	virtual void addSamples() = 0;
	virtual bool terminationCheck() { return false; }
	virtual void downloadSamples() = 0;
	virtual unsigned int getNumGeneratedSamples() = 0;

	void setUseMultipleBlocks(bool value) { useMultipleBlocks_ = value; }
	bool getUseMultipleBlocks() const { return useMultipleBlocks_; }

	void setValidateResults(bool value) { validateResults_ = value; }
	bool getValidateResults() const { return validateResults_; }
protected:
	bool useMultipleBlocks_;
	bool validateResults_;

	curandState* d_randStates_;
	GPUDart::GPUInterval* d_intervals_;
	GPUMDVector d_samples_;
	GPUMDVector d_refPoints_;

	unsigned int h_dimensions_;
	unsigned int h_numSamples_;
	unsigned int h_failureCount_;
	unsigned int h_m_;
	DATA_TYPE h_radius_;

	unsigned int* h_buffer_;
};

__inline__ __device__ void GPUDart::GPUInterval::insert(short pos, DATA_TYPE begin, DATA_TYPE end)
{
	if (numUsed == MAX_INTERVAL)
	{
		printf("interval overflow!!!\n");
		return;
	}

	short prevIndex = prev[pos];

	short newIndex = numUsed++;
	intervalBegins[newIndex] = begin;
	intervalEnds[newIndex] = end;
	prev[newIndex] = prevIndex;
	next[newIndex] = pos;
	prev[pos] = newIndex;
	if (prevIndex != -1)
	{
		next[prevIndex] = newIndex;
	}
	else
		head = newIndex;

	++numIntervals;
}

__inline__ __device__ void GPUDart::GPUInterval::erase(short from, short to)
{
	for (short it = from; it != to; it = next[it])
	{
		--numIntervals;
	}

	short prevIndex = prev[from];
	prev[to] = prevIndex;
	if (prevIndex != -1)
	{
		next[prevIndex] = to;
	}
	else
		head = to;
}

__inline__ __device__ short GPUDart::GPUInterval::advance(short* indexArray, short it, short step)
{
	for (int i = 0; i < step; ++i)
		it = indexArray[it];
	return it;
}

__inline__ __device__ short GPUDart::GPUInterval::lower_bound(DATA_TYPE* dataArray, short* indexArray, short first, short last, DATA_TYPE val, short* distance)
{
	short it;
	short count, step;
	count = *distance;
	while (count > 0)
	{
		step = count / 2; 
		it = advance(indexArray, first, step);
		if (dataArray[it] < val) {
			first = indexArray[it];
			count -= step + 1;
			*distance -= step + 1;
		}
		else 
			count = step;
	}
	return first;
};

__inline__ __device__ short GPUDart::GPUInterval::upper_bound(DATA_TYPE* dataArray, short* indexArray, short first, short last, DATA_TYPE val, short* distance)
{
	short it;
	short count, step;
	count = *distance;
	while (count > 0)
	{
		step = count / 2; 
		it = advance(indexArray, first, step);
		if (dataArray[it] <= val) {
			first = indexArray[it];
			count -= step + 1;
			*distance -= step + 1;
		}
		else 
			count = step;
	}
	return first;
};

#endif